//
//  Convolutioner_TimeDomain_OpenCL.cpp
//  FilterModule_Standalone
//
//  Created by kadlubek47 on 06.08.2013.
//  Copyright (c) 2013 kadlubek47. All rights reserved.
//

#include "Convolutioner_TimeDomain_OpenCL.h"



const char Convolutioner_TimeDomain_OpenCL::kernelPath_[] = "convolute_kernel_time_domain.cl";



#pragma mark Convolutioner Creation

Convolutioner_TimeDomain_OpenCL::Convolutioner_TimeDomain_OpenCL(Parameters_Convolutioner::ProcessingMethodID processingMethodID, const IRF* irf): OpenCLProcessorBase((char*)kernelPath_) {
    //must be creation
    create_Kernels();
    create_BufferGlobalParameters();
    //.
    
    irLength_ = irf->N_;
    recreateAndSet_BufferIRF(irf);
    toSet_BufferGlobalParameters = true;
}

Convolutioner_TimeDomain_OpenCL::~Convolutioner_TimeDomain_OpenCL() {
    if(bufferIn_created_)
        clReleaseMemObject(bufferIn_);
    if(bufferOut_created_)
        clReleaseMemObject(bufferOut_);
    if(bufferIR_created_)
        clReleaseMemObject(bufferIR_);
    clReleaseMemObject(bufferGlobalParameters_);
    
    clReleaseKernel(kernel_convolution_);
}

void Convolutioner_TimeDomain_OpenCL::create_Kernels() {
    kernel_convolution_ = clCreateKernel(program_, "convoluteChannels", &lastCommandStatus_);
}

void Convolutioner_TimeDomain_OpenCL::create_BufferGlobalParameters() {
    bufferGlobalParameters_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, 2 * sizeof(cl_uint), NULL, &lastCommandStatus_);
    lastCommandStatus_ |= clSetKernelArg(kernel_convolution_, 3/*KERNEL NUMBER ARGUMENT*/, sizeof(cl_mem), &bufferGlobalParameters_ );
}



#pragma mark main_function

void Convolutioner_TimeDomain_OpenCL::convolute(AudioBlocks<float_type>& audio) {
    if(blockSize_ != audio.numOfSamples_)
        toSet_BufferGlobalParameters = true;
    blockSize_ = audio.numOfSamples_;
    numOfChannels_ = audio.numOfChannels_;
    
    //Update channelsWindow
    channelsWindowLinear_.update(audio, irLength_-1);
    //.
    
    //Check if must be changed input output buffers length
    if(channelsWindowLinear_.getLength() != bufferIn_length_)
        recreate_BufferIn();
    if(numOfChannels_* blockSize_ != bufferOut_length_)
        recreate_BufferOut();
    
    if(toSet_BufferGlobalParameters) {
        set_BufferGlobalParameters();
        toSet_BufferGlobalParameters = false;
    }
    //.
    
    //Process and convert from ChannelsWindowLinear to AudioBlocks
    
    float_type* outputChannels = new float_type[audio.numOfSamples_*audio.numOfChannels_];
    loadBlocksOfChannelsIntoGPU(channelsWindowLinear_.getWindowLinear());
    executeKernel();
    getBlocksOfChannelsFromGPU(outputChannels);
    FromChannelsWindowsLinear_ToAudioBlocks_Converter::convert(outputChannels, audio);
    delete [] outputChannels;
    //. 
}

void Convolutioner_TimeDomain_OpenCL::setNewIR(const IRF* irf) {
    toSet_BufferGlobalParameters = true;
    irLength_ = irf->N_;
    recreateAndSet_BufferIRF(irf);
}

void Convolutioner_TimeDomain_OpenCL::setProcessingMethod (Parameters_Convolutioner::ProcessingMethodID& param) {
    
}



#pragma mark subordinate_functions

void Convolutioner_TimeDomain_OpenCL::recreateAndSet_BufferIRF(const IRF* irf) {
    //Check if must be recreated
    if(bufferIR_length_ != irf->N_) {
        if(bufferIR_created_)
            clReleaseMemObject(bufferIR_);
        bufferIR_length_ = irf->N_;
        bufferIR_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, bufferIR_length_ * sizeof(float_type), NULL, &lastCommandStatus_);
        bufferIR_created_ = true;
    }
    //.
    
    //transfer data
    lastCommandStatus_ = clEnqueueWriteBuffer(cmdQueue_, bufferIR_, CL_FALSE, 0, bufferIR_length_ * sizeof(float), irf->h_, 0, NULL, NULL);
    //.
    
    //set argument
    lastCommandStatus_  = clSetKernelArg(kernel_convolution_, 2, sizeof(cl_mem), &bufferIR_);
    //.
}

void Convolutioner_TimeDomain_OpenCL::recreate_BufferIn() {
    //Delete previous buffers with different size
    if( bufferIn_created_ ) {
        clReleaseMemObject(bufferIn_);
    }
    //.
    
    //Create
    bufferIn_length_ = channelsWindowLinear_.getLength();
    bufferIn_   = clCreateBuffer(context_, CL_MEM_READ_ONLY,  bufferIn_length_ * sizeof(float), NULL, &lastCommandStatus_);
    //.
    
    //Set kernel arguments
    lastCommandStatus_  = clSetKernelArg(kernel_convolution_, 0, sizeof(cl_mem), &bufferIn_);
    //.
    
    //Set flag of creation
    bufferIn_created_ = true;
    //.
    
}

void Convolutioner_TimeDomain_OpenCL::recreate_BufferOut() {
    //Delete if existed
    if( bufferOut_created_ ) {
        clReleaseMemObject(bufferOut_);
    }
    //.

    //Create
    bufferOut_length_ = numOfChannels_*blockSize_;
    bufferOut_  = clCreateBuffer(context_, CL_MEM_WRITE_ONLY, bufferOut_length_ * sizeof(float), NULL, &lastCommandStatus_);
    //.
    
    //Set kernel arguments
    lastCommandStatus_  = clSetKernelArg(kernel_convolution_, 1, sizeof(cl_mem), &bufferOut_);
    //.
    
    //Set flag of creation
    bufferOut_created_ = true;
    //.
}

void Convolutioner_TimeDomain_OpenCL::set_BufferGlobalParameters() {
    cl_uint* info = new cl_uint[2];
    info[0] = (cl_uint)blockSize_;      //< info[0] == channelLength_   Input Block Size(latency;)
    info[1] = (cl_uint)irLength_;           //< info[1] == irLength_        Impulse Response Length
    lastCommandStatus_ = clEnqueueWriteBuffer(cmdQueue_, bufferGlobalParameters_, CL_FALSE, 0, sizeof(cl_uint) * 2, info, 0, NULL, NULL);
}

void Convolutioner_TimeDomain_OpenCL::loadBlocksOfChannelsIntoGPU(float* channels) {
    lastCommandStatus_ = clEnqueueWriteBuffer(cmdQueue_, bufferIn_, CL_FALSE, 0, bufferIn_length_ * sizeof(float), channels, 0, NULL, NULL);
}

void Convolutioner_TimeDomain_OpenCL::executeKernel() {
    size_t globalWorkSize[1];
    globalWorkSize[0] = bufferOut_length_;
    lastCommandStatus_ = clEnqueueNDRangeKernel(cmdQueue_, kernel_convolution_, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL);
}

void Convolutioner_TimeDomain_OpenCL::getBlocksOfChannelsFromGPU (float* destiny) {
    lastCommandStatus_ = clEnqueueReadBuffer(cmdQueue_, bufferOut_, CL_FALSE, 0, bufferOut_length_ * sizeof(float), destiny, 0, NULL, NULL);
    clFlush(cmdQueue_);
    clFinish(cmdQueue_);
}